
#include <stdio.h>
#define TAMANHO 5

__global__ void MultiplicaVetor (float4* vetor, float* acumulador)
{
	//__shared__ float acumulador = 0.0f;
	
	if(threadIdx.x < TAMANHO)
	{
		vetor[threadIdx.x].x *= 2.0f;
		//vetor[blockIdx.x].x = 33.3f;
		vetor[threadIdx.x].y *= 2.0f;
		vetor[threadIdx.x].z *= 2.0f;
		vetor[threadIdx.x].w *= 2.0f;
	}

	__syncthreads();
		
	if(threadIdx.x == 0)
	{
		float acumulador = 0.0f;
		for(int i=0; i<TAMANHO; i++)
		{
			acumulador += vetor[i].x;
		}
	}
}



int main(void)
{

	float4* vetor;
	float4* d_vetor;
	float acumulador = 0.0f;
	float d_acumulador;
	
	vetor = (float4*)malloc(sizeof(float4)*TAMANHO);
	
	for(int i=0; i<TAMANHO; i++)
	{
		vetor[i].x = i;
		vetor[i].y = i;
		vetor[i].z = i;
		vetor[i].w = i;
	}
	
	printf("\n:::: Vetor antes da chamada do Kernel ::::\n");
	printf("==> Acumulador: %f \n", acumulador);
	for(int i =0; i<TAMANHO; i++)
	{
		printf(":::: [%d] >> vetor[%d].x = %f :: vetor[%d].y = %f :: vetor[%d].z = %f :: vetor[%d].w = %f \n", i, i, vetor[i].x, i, vetor[i].y, i, vetor[i].z, i, vetor[i].w);
	}
	printf("\n");
	
	cudaMalloc( (void**)&d_vetor,      sizeof(float4)*TAMANHO );
	cudaMalloc( &d_acumulador, sizeof(float)          );
	
	cudaMemcpy( d_vetor,      vetor,      sizeof(float4)*TAMANHO, cudaMemcpyHostToDevice );
	cudaMemcpy( d_acumulador, acumulador, sizeof(float),          cudaMemcpyHostToDevice );
	
	// kernel !!!
	MultiplicaVetor<<< 1, TAMANHO >>>(d_vetor, d_acumulador);

	cudaMemcpy( vetor,      d_vetor,      sizeof(float4)*TAMANHO,  cudaMemcpyDeviceToHost );
	cudaMemcpy( acumulador, d_acumulador, sizeof(float),           cudaMemcpyDeviceToHost );

	printf("\n:::: Vetor DEPOIS da chamada do Kernel ::::\n");
	printf("==> Acumulador: %f \n", acumulador);
	for(int i =0; i<TAMANHO; i++)
	{
		printf(":::: [%d] >> vetor[%d].x = %f :: vetor[%d].y = %f :: vetor[%d].z = %f :: vetor[%d].w = %f \n", i, i, vetor[i].x, i, vetor[i].y, i, vetor[i].z, i, vetor[i].w);
	}
	printf("\n");

	free(vetor);
	cudaFree(d_vetor);
	
}